/*
 * @Author: xiangru.xiao xiangru.xiao@mthreads.com
 * @Date: 2023-05-23 20:04:37
 * @LastEditors: xiangru.xiao xiangru.xiao@mthreads.com
 * @LastEditTime: 2023-05-25 10:59:28
 * Copyright (c) 2023 by Moore Threads, All Rights Reserved. 
 */

#include <string>
#include <vector>
#include "objectfactory.h"
#include "node.h"
#include "cu_helper.h"
#include "kernel.cuh"
#define EPS 1e-6
#define TEST_ITERS 10
struct params
{
    int length;
    int batch;
};
bool testResult(cufftHandle plan, float* h_x, void* d_x, int length, int batch);

class FFT_test : public Node
{
public:
    FFT_test()
    {
        nodeUint = "GBPS";
    }
    std::string name() override
    {
        return "fft-test";
    }
    std::string describe() override
    {
        return "快速傅里叶变换";
    }
    Node::NodeType  nodeType() override
    {
        return NodeTP_MemRw;
    }
protected:
    bool startup() override;
    bool perform() override;
    void shutdown() override;
private:
    std::vector<params> paramList() const;
private:
    int        deviceid;
};

bool FFT_test::startup()
{
    deviceid = deviceinfo->deviceID();
    CHECK_CUDA_ERROR(cudaSetDevice(deviceid));
    return true;
}

bool FFT_test::perform()
{
    float totalValue = 0.0;
    std::vector<params> ls = paramList();
    char* strNum = std::getenv("FFT_TEST_LEVEL");
    int caseNum = ls.size();
    if(strNum)
    {
        caseNum = std::min(caseNum, atoi(strNum));
    }
    for (int i = 0; i < caseNum; ++i) {
        int length = ls[i].length;
        int batch = ls[i].batch;
        const int complex_bytes = sizeof(float) * 2 * length * batch;
        void *h_x, *d_x;
        cufftHandle plan;
        CHECK_CUDA_ERROR(cudaMallocHost(&h_x, complex_bytes));
        CHECK_CUDA_ERROR(cudaMalloc(&d_x, complex_bytes));
        CHECK_CUDA_ERROR(cufftPlan1d(&plan, length, CUFFT_C2C, batch));
        if(testResult(plan, (float*)h_x, d_x, length, batch))
        { //result correct and then do profile
            float totalTime  = 0;
            {
                CUDAEvent evt(&totalTime);
                for (int j = 0; j < TEST_ITERS; ++j) {
                      CHECK_CUDA_ERROR(cufftExecC2C(plan, (cufftComplex*)d_x, 
                                        (cufftComplex*)d_x, CUFFT_FORWARD));
                      CHECK_CUDA_ERROR(cudaDeviceSynchronize());
                }
            }
            CHECK_CUDA_ERROR(cufftDestroy(plan));
            CHECK_CUDA_ERROR(cudaFreeHost(h_x));
            CHECK_CUDA_ERROR(cudaFree(d_x));
            totalValue += 2 * float(complex_bytes) * TEST_ITERS / totalTime / 1e6f;
        }
        else
        { //get wrong result
            nodeValue = 0.0;
            CHECK_CUDA_ERROR(cufftDestroy(plan));
            CHECK_CUDA_ERROR(cudaFreeHost(h_x));
            CHECK_CUDA_ERROR(cudaFree(d_x));
            return false;
        }
        
    }
    nodeValue = totalValue;
    return true;
}

void FFT_test::shutdown()
{

}

std::vector<params> FFT_test::paramList() const
{
    std::vector<params> list;
    params params_[] = {{128, 1<<18},
                        {256, 1<<17},
                        {512, 1<<16},
                        {1024, 1<<15},
                        {2048, 1<<14},
                        };
    const size_t nSizes =  sizeof(params_) / sizeof (params);
    for (int i = 0; i < nSizes; ++i) {
        list.push_back(params_[i]);
    }
    return list;
}

REGISTER_OBJECT(Node, FFT_test)

bool testResult(cufftHandle plan, float* h_x, void* d_x, int length, int batch)
{
    int N = length * std::min(batch, 32);
    float* h_out = (float*)malloc(sizeof(float) * 2 * N);
    for(int i = 0; i < N; i++)
    {
        h_x[2 * i] = rand();
        h_x[2 * i + 1] = 0;
    }
    CHECK_CUDA_ERROR(cudaMemcpy(d_x, h_x, sizeof(float) * 2 * N, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERROR(cufftExecC2C(plan, (cufftComplex*)d_x, 
                                        (cufftComplex*)d_x, CUFFT_FORWARD));
    CHECK_CUDA_ERROR(cudaDeviceSynchronize());
    CHECK_CUDA_ERROR(cufftExecC2C(plan, (cufftComplex*)d_x, 
                                        (cufftComplex*)d_x, CUFFT_INVERSE));
    CHECK_CUDA_ERROR(cudaDeviceSynchronize());
    CHECK_CUDA_ERROR(cudaMemcpy(h_out, d_x, sizeof(float) * 2 * N, cudaMemcpyDeviceToHost));

    double maxv = 0;
    double nrmse = 0;  // normalized root mean square error
    for (int i = 0; i < N; i++) {
      double dr = h_x[2 * i] - h_out[2 * i]/length;
      double di = h_x[2 * i + 1] - h_out[2 * i + 1]/length;
      maxv = fabs(h_x[2 * i]) > maxv ? fabs(h_x[2 * i]) : maxv;
      maxv = fabs(h_x[2 * i + 1]) > maxv ? fabs(h_x[2 * i + 1]) : maxv;
      nrmse += ((dr * dr) + (di * di));
    }
    nrmse /= (double)(N);
    nrmse = sqrt(nrmse);
    nrmse /= maxv;
    
    free(h_out);
    return nrmse < EPS;
}